#include <stdio.h> // printf
#include <assert.h>
#include <cuda.h> // CUDA library
//#include <helper_cuda.h>


const unsigned int SIZE=1;

//__global__ void inc(volatile float *heat) {
__global__ void inc(float *heat) {

	//version con condition race
	//*heat = *heat+1.0f;
	// version sin condition race
	atomicAdd(heat, 1.0f);
}



int main(void)
{

	float *h_heat = NULL; // pointer to a float in host memory
	float *d_heat = NULL; // pointer to a float in device memory
	
	h_heat = (float *) calloc(SIZE, sizeof(float));
	assert(h_heat);

	// cutilSafeCall isn't present in cuda 5
	//cutilSafeCall(cudaMalloc(&d_heat, SIZE*sizeof(float)));
	//checkCudaErrors( cudaMalloc(&d_heat, SIZE*sizeof(float)));
 	cudaMalloc(&d_heat, SIZE*sizeof(float));

	*h_heat = 42.0f;
	// cutilSafeCall isn't present in cuda 5
	//cutilSafeCall(cudaMemcpy(d_heat, h_heat, SIZE*sizeof(float),cudaMemcpyDefault));
 	//checkCudaErrors(cudaMemcpy(d_heat, h_heat, SIZE*sizeof(float),cudaMemcpyDefault));
 	cudaMemcpy(d_heat, h_heat, SIZE*sizeof(float),cudaMemcpyDefault);
	
	*h_heat = 0.0f;
	// cutilSafeCall isn't present in cuda 5
	//cutilSafeCall(cudaMemcpy(h_heat, d_heat, SIZE*sizeof(float),cudaMemcpyDefault));
 	//checkCudaErrors(cudaMemcpy(h_heat, h_heat, SIZE*sizeof(float),cudaMemcpyDefault));
 	cudaMemcpy(h_heat,d_heat, SIZE*sizeof(float),cudaMemcpyDefault);
	

	inc<<<32768,512>>>(d_heat);
	//cutilCheckMsg("inc kernel failed");
	//cutilSafeCall(cudaDeviceSynchronize()); 


 	cudaMemcpy(h_heat,d_heat, SIZE*sizeof(float),cudaMemcpyDefault);

	cudaDeviceSynchronize(); 
 




	printf("Heat: %f\n", *h_heat);
	// cutilSafeCall isn't present in cuda 5
	//cutilSafeCall(cudaFree(d_heat));
 	//checkCudaErrors(cudaFree(d_heat));
 	//cudaFree(d_heat);
	
	free(h_heat);
	return 0;
}





//nvcc -O3 -arch=sm_20 -I/opt/cudasdk/C/common/inc -o sample.o -c sample1.cu
//nvcc -O3 -arch=sm_20 -I/opt/cudasdk/C/common/inc -o sample sample.o
